Skip to content

Triton monolithic head-loss kernel (#507)#558

Merged
jlamypoirier merged 6 commits into
jlp_monolithic_head_lossfrom
jlp_monolithic_triton
Jul 14, 2026
Merged

Triton monolithic head-loss kernel (#507)#558
jlamypoirier merged 6 commits into
jlp_monolithic_head_lossfrom
jlp_monolithic_triton

Conversation

@jlamypoirier

@jlamypoirier jlamypoirier commented Jul 8, 2026

Copy link
Copy Markdown
Collaborator

Claude Opus 4.8 (1M context) note: authored this PR.

Follow-up to #549, stacked on jlp_monolithic_head_loss. Gives the monolithic loss a single Triton kernel for its shared-softmax path, alongside the existing torch.compile path.

What it does

The monolithic loss can now run its shared softmax on one @triton_jit kernel via use_triton, covering the label-based objective set: cross-entropy (label), z_loss, the GRPO objective, and the GSPO objective. Any subset sharing one softmax works (CE+z, GRPO+z, GSPO+z, …), at most one of each kind.

Distillation has no Triton fused kernel yet — it is rejected at config time when use_triton=True, and falls back to the compiled path when use_triton is left at its None (auto) default. DPO stays separate as before.

Design

Every enabled loss stores its own forward scalar, but the gradients superpose into two per-row coefficients:

grad_j = prob_coeff · softmax_j − label_coeff · δ_{j, label}

so the gradient-write loop is loss-agnostic and each loss's math only contributes to the coefficients (z-loss adds to prob_coeff only; CE/GRPO/GSPO share the same label column). The kernel is gated purely by the presence of each loss's input/output pointers — no separate enable flags.

Three dispatch cases mirror the existing per-loss Triton kernels:

  • No tensor parallelism, no GSPO — one fused kernel, softmax computed in-kernel.
  • Tensor parallelism — a reduced forward pass (reusing the shared cross-entropy forward kernel) feeds a load-mode kernel.
  • GSPO present — its eager segment seam (from Fuse the GSPO segment seam into compiled blocks (#507) #554) runs between the forward and the load-mode backward, handing the kernel a per-token coefficient that superposes with the in-kernel losses.

GSPO's seam-argument assembly is factored out of finish so the compiled and Triton paths share it.

Config

use_triton: bool | None on MonolithicLossConfig:

  • True — rejects unsupported children (distillation, duplicate kinds) at validate time.
  • None (default) — Triton on GPU only if every child is supported, else the compiled monolithic.
  • False — always compiled.

Validation

Validated on an H100 GPU (FAST_LLM_SKIP_TRITON_AUTOTUNE=TRUE):

  • test_lm_losses functional — the monolithic kernel against the PyTorch and standalone-Triton references (rms-relative ~1e-7): CE / z / CE+z, plus the GRPO and GSPO objectives through their kernel slots, across masking, logit/grad scaling, fp16, gradient accumulation, and the looped multi-block path.
  • A metrics-without-gradient case — a zero-weight policy loss still logging its diagnostics (loss / new-log-probs / entropy) with no gradient written.
  • test_lm_head fused_triton configs — single-pass and the GSPO seam, across masking / splits / bf16 / logit scaling / softcap.
  • Distributed gloo world_size=2 monolithic subtests for the tensor-parallel path (CE+z, GRPO, GSPO).
  • A plain-CPU run leaves the compiled path intact and skips the Triton head configs.

The Triton interpreter (TRITON_INTERPRET=1) can't run this kernel — the grid trips an InterpreterError on a constexpr-shaped dimension (reproduces on main, so it's pre-existing) — so validation of the monolithic Triton path is GPU-only.

Follow-ups: distillation in the Triton kernel, and the GRPO/GSPO metrics rework.

🤖 Generated with Claude Code

jlamypoirier added a commit that referenced this pull request Jul 10, 2026
Bring main (via the updated compiled base) into the triton PR. The metrics
rework and triton work merge cleanly; fix the leftover references to the
renamed metrics enum (`GRPOMetricsLevel` -> `PolicyMetricsLevel`) in the triton
child classification (`config.py`, `monolithic.py`). GSPO metrics are already
rejected for any monolithic composite, so the triton path only rejects GRPO
metrics (which it doesn't compute).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
@jlamypoirier
jlamypoirier force-pushed the jlp_monolithic_triton branch from 73683f7 to 284fe2f Compare July 10, 2026 20:03
@jlamypoirier
jlamypoirier changed the base branch from jlp_monolithic_head_loss to jlp_monolithic_metrics July 10, 2026 20:04
@jlamypoirier
jlamypoirier force-pushed the jlp_monolithic_metrics branch from 32c3943 to 2368650 Compare July 11, 2026 01:12
@jlamypoirier
jlamypoirier force-pushed the jlp_monolithic_triton branch from 284fe2f to 8c9f052 Compare July 11, 2026 01:12
@jlamypoirier
jlamypoirier force-pushed the jlp_monolithic_metrics branch from 2368650 to 208d509 Compare July 13, 2026 15:53
@jlamypoirier
jlamypoirier force-pushed the jlp_monolithic_triton branch from 8c9f052 to f3771ce Compare July 13, 2026 15:59
Base automatically changed from jlp_monolithic_metrics to jlp_monolithic_head_loss July 13, 2026 16:30
Add the triton implementation of the monolithic shared-softmax loss and its
policy-metric support, on top of the compiled path:

- The triton kernel fuses the label-based objective set (CE, z-loss, GRPO, GSPO)
  over one softmax, superposing their gradients.
- Policy metrics: the kernel accumulates the entropy's `Σ exp·logits_norm` for
  free in the backward re-stream; the layer reduces the family eagerly. Metrics
  force the reduced forward pass so the softmax is available in every case.
- `MonolithicLossConfig.use_triton` selects the path (default: use if available),
  with per-kind fusion limits validated at config time.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
@jlamypoirier
jlamypoirier force-pushed the jlp_monolithic_triton branch from f3771ce to 2e9e206 Compare July 13, 2026 16:38
…riton path (#507)

- `compute_gspo_metrics`: fuse softmax + predicted-logits + entropy into one compiled
  pass (new `_gspo_metrics_forward`), leaving only the segment `index_add_` eager. It was
  fully eager, computing entropy as a separate reduction over a vocab-wide temporary —
  mirrors the compiled `compute_grpo_metrics`.
- GSPO `fused_core`: reduce `entropy_per_token` inside the compiled boundary (only when
  metrics are requested) and hand it to `finish`, rather than holding `logits_norm` (a
  vocab-wide tensor) across the eager seam. Fixes the monolithic GSPO single-loss peak
  memory and the fused-metrics speed.
- Monolithic triton kernel: skip the softmax and the backward `exp` for a masked row when
  only a single label-based loss is active (nothing else needs the softmax on a label-less
  row). Folded into the existing backward loop rather than an early `return`, which triton
  compiles into heavy register spills in this large kernel. Brings single-loss cross-entropy
  and GRPO to parity with the dedicated kernels.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
@jlamypoirier
jlamypoirier force-pushed the jlp_monolithic_triton branch from c66f242 to b64da2b Compare July 13, 2026 23:32
jlamypoirier and others added 4 commits July 14, 2026 14:31
Route the triton monolithic path through per-loss methods instead of a
central kind classification, mirroring the compiled `_monolithic_core`
loop. Each combinable loss owns `triton_add_inputs` / `triton_seam` /
`triton_finish` and a `triton_kind` on its config (the single source of
truth consulted by both config validation and the runtime); the driver
iterates its children filling a shared `_TritonContext`, with no
per-kind branching or duplicated classification. The `@triton.jit`
kernel keeps its fixed per-slot signature — the one boundary where the
loss kinds are still named.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The ce+z monolithic triton tests never passed `grpo=` / `gspo_coeff=`, so
those kernel slots were covered only by the full head tests. Add
`_test_monolithic_grpo_loss_triton` and `_test_monolithic_gspo_loss_triton`,
each checking a single policy loss through the monolithic kernel against the
standalone python reference (GSPO's per-token coefficient produced by the
eager segment seam, as the driver runs it), plus matching tensor-parallel
subtests in the distributed harness.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The metrics' `Σ exp·logits_norm` is accumulated in the backward re-stream,
which previously ran only when a gradient was requested — so a fused
policy loss with metrics but no gradient (e.g. a zero-weight child logging
diagnostics) crashed the triton path while the compiled path handled it.
Run the re-stream when either a gradient or the metrics are requested, and
gate only the gradient store on `grad_logits_ptr`, so the two backends
agree.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
predicted_logit is already initialized to 0.0, so the else-branch reassigned
the same value. No behavioral change.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
@jlamypoirier
jlamypoirier merged commit 52fabf0 into jlp_monolithic_head_loss Jul 14, 2026
1 check passed
@jlamypoirier
jlamypoirier deleted the jlp_monolithic_triton branch July 14, 2026 20:53
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant